#include "kernel_operator.h"

using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;



class KernelDawsn
{
public:
    __aicore__ inline KernelDawsn() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,
                            uint32_t ALIGN_NUM, uint32_t core_size, 
                            uint32_t core_remain, uint32_t ub_data_num)
    {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        uint32_t blockIdx = GetBlockIdx();

        this->coreDataNum = core_size + (GetBlockNum() == GetBlockIdx() + 1 ? core_remain : 0);
        this->tileDataNum = ub_data_num;

        auto startPointer = core_size * GetBlockIdx();
        xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + startPointer, this->coreDataNum);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + startPointer, this->coreDataNum);

        this->tileNum = this->coreDataNum / this->tileDataNum + (this->coreDataNum % this->tileDataNum > 0);

        pipe.InitBuffer(xQueue, BUFFER_NUM, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(yQueue, BUFFER_NUM, this->tileDataNum * sizeof(DTYPE_Y));
        pipe.InitBuffer(All0, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(All1, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(t1Buf, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(t2Buf, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(t3Buf, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(t4Buf, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(t5Buf, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(rBuf, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(Sign, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(Bits, this->tileDataNum * sizeof(uint8_t));
        pipe.InitBuffer(Mask, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(Mask2, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(All, this->tileDataNum * sizeof(DTYPE_X));





    }

    __aicore__ inline void Process()
    {
        uint32_t loopCount = this->tileNum;
        for (int32_t i = 0; i < loopCount - 1; i++) 
        {
            CopyIn(i, this->tileDataNum);
            Compute(i, this->tileDataNum);
            CopyOut(i, this->tileDataNum);
        }

        uint32_t length = this->coreDataNum - this->tileDataNum * (loopCount - 1);
        CopyIn(loopCount - 1, length);
        Compute(loopCount - 1, length);
        CopyOut(loopCount - 1, length);
    }

private:

    __aicore__ inline void CopyIn(int32_t progress, uint32_t length) 
    {
        LocalTensor<DTYPE_X> xLocal = xQueue.AllocTensor<DTYPE_X>();

        DataCopy(xLocal, xGm[progress * this->tileDataNum], length);
        
        xQueue.EnQue(xLocal);
    }

    __aicore__ inline void Compute(int32_t progress, uint32_t length)
    {
        LocalTensor<DTYPE_X> xLocal = xQueue.DeQue<DTYPE_X>();
        LocalTensor<DTYPE_Y> yLocal = yQueue.AllocTensor<DTYPE_Y>();
        LocalTensor<DTYPE_X> all0 = All0.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> all1 = All1.Get<DTYPE_X>();
        LocalTensor<uint8_t> bits = Bits.Get<uint8_t>();
        LocalTensor<DTYPE_X> sign = Sign.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> mask = Mask.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> mask2 = Mask2.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> all = All.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> rLocal = rBuf.Get<DTYPE_X>();

        // Duplicate(all1, (DTYPE_X)1.0, length);
        Muls(all1, all1, (DTYPE_X)0.0, length);
        Adds(all1, all1, (DTYPE_X)1.0, length);
        Muls(all0, all0, (DTYPE_X)0.0, length);
        Muls(rLocal, rLocal, (DTYPE_X)0.0, length);


        Compare(bits, xLocal, all0, CMPMODE::GT, length);     //1 1 0 0 1 0 1
        Select(sign, bits, all1, (DTYPE_X)-1.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        Abs(xLocal, xLocal, length);

        Muls(all, all, (DTYPE_X)0.0, length);
        Adds(all, all, (DTYPE_X)3.25, length);
        Compare(bits, xLocal, all, CMPMODE::LT, length);
        Select(mask, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        _calc_condition_lt_three_p_two_five(xLocal, yLocal, length);
        Mul(yLocal, yLocal, mask, length);
        Add(rLocal, rLocal, yLocal, length);

        Muls(all, all, (DTYPE_X)0.0, length);
        Adds(all, all, (DTYPE_X)3.25, length);
        Compare(bits, xLocal, all, CMPMODE::GE, length);
        Select(mask, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        Muls(all, all, (DTYPE_X)0.0, length);
        Adds(all, all, (DTYPE_X)6.25, length);
        Compare(bits, xLocal, all, CMPMODE::LT, length);
        Select(mask2, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        Add(mask, mask, mask2, length);
        Muls(all, all, (DTYPE_X)0.0, length);
        Adds(all, all, (DTYPE_X)2.0, length);
        Compare(bits, mask, all, CMPMODE::EQ, length);
        Select(mask2, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        _calc_condition_lt_six_p_two_five(xLocal, yLocal, length);
        Mul(yLocal, yLocal, mask2, length);
        Add(rLocal, rLocal, yLocal, length);


        Muls(all, all, (DTYPE_X)0.0, length);
        Adds(all, all, (DTYPE_X)6.25, length);
        Compare(bits, xLocal, all, CMPMODE::GE, length);
        Select(mask, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        Muls(all, all, (DTYPE_X)0.0, length);
        Adds(all, all, (DTYPE_X)1e9, length);
        Compare(bits, xLocal, all, CMPMODE::LE, length);
        Select(mask2, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        Add(mask, mask, mask2, length);
        Muls(all, all, (DTYPE_X)0.0, length);
        Adds(all, all, (DTYPE_X)2.0, length);
        Compare(bits, mask, all, CMPMODE::GE, length);
        Select(mask2, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        _calc_condition_le_one_e_nine(xLocal, yLocal, length);
        Mul(yLocal, yLocal, mask2, length);
        Add(rLocal, rLocal, yLocal, length);

        Muls(all, all, (DTYPE_X)0.0, length);
        Adds(all, all, (DTYPE_X)1e9, length);
        Compare(bits, xLocal, all, CMPMODE::GT, length);
        Select(mask, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
        _calc_condition_gt_one_e_nine(xLocal, yLocal, length);
        Mul(yLocal, yLocal, mask, length);
        Add(rLocal, rLocal, yLocal, length);
        
        Mul(yLocal, rLocal, sign, length);

        yQueue.EnQue(yLocal);
        xQueue.FreeTensor(xLocal);
    }

    __aicore__ inline void CopyOut(int32_t progress, uint32_t length) {
        LocalTensor<DTYPE_Y> yLocal = yQueue.DeQue<DTYPE_Y>();
        DataCopy(yGm[progress * this->tileDataNum], yLocal, length);
        yQueue.FreeTensor(yLocal);
    }

    __aicore__ inline void _calc_condition_lt_three_p_two_five(LocalTensor<DTYPE_X> &xLocal, 
                                                                LocalTensor<DTYPE_Y> &yLocal,
                                                                uint32_t length)
    {
        LocalTensor<DTYPE_X> t1Local = t1Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t2Local = t2Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t3Local = t3Buf.Get<DTYPE_X>();


        Mul(t1Local, xLocal, xLocal, length);

        Muls(t2Local, t2Local, (DTYPE_X)0.0, length);
        Adds(t2Local, t2Local, (DTYPE_X)1.13681498971755972054E-11, length);

        {
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)8.49262267667473811108E-10, length);
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)1.94434204175553054283E-8, length);
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)9.53151741254484363489E-7, length);
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)3.07828309874913200438E-6, length);
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)3.52513368520288738649E-4, length);
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)-8.50149846724410912031E-4, length);
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)4.22618223005546594270E-2, length);
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)-9.17480371773452345351E-2, length);
            Mul(t2Local, t1Local, t2Local, length);
            Adds(t2Local, t2Local, (DTYPE_X)9.99999999999999994612E-1, length);
        }

        Muls(t3Local, t3Local, (DTYPE_X)0.0, length);
        Adds(t3Local, t3Local, (DTYPE_X)1.13681498971755972054E-11, length);
        
        {
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)1.48864681368493396752E-9, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)5.21265281010541664570E-8, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)1.27258478273186970203E-6, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)2.32490249820789513991E-5, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)3.25524741826057911661E-4, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)3.48805814657162590916E-3, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)2.79448531198828973716E-2, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)1.58874241960120565368E-1, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)5.74918629489320327824E-1, length);
            Mul(t3Local, t1Local, t3Local, length);
            Adds(t3Local, t3Local, (DTYPE_X)1.00000000000000000539E0, length);
        }

        Mul(yLocal, xLocal, t2Local, length);
        Div(yLocal, yLocal, t3Local, length);
    }

    __aicore__ inline void _calc_condition_lt_six_p_two_five(LocalTensor<DTYPE_X> &xLocal, 
                                                                LocalTensor<DTYPE_Y> &yLocal,
                                                                uint32_t length)
    {
        LocalTensor<DTYPE_X> t1Local = t1Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t2Local = t2Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t3Local = t3Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t4Local = t4Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t5Local = t5Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> all1 = All1.Get<DTYPE_X>();

        Mul(t1Local, xLocal, xLocal, length);
        Div(t2Local, all1, t1Local, length);
        Div(t3Local, all1, xLocal, length);


        Muls(t4Local, t4Local, (DTYPE_X)0.0, length);
        Adds(t4Local, t4Local, (DTYPE_X)5.08955156417900903354E-1, length);
        {
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)-2.44754418142697847934E-1, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)9.41512335303534411857E-2, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)-2.18711255142039025206E-2, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)3.66207612329569181322E-3, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)-4.23209114460388756528E-4, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)3.59641304793896631888E-5, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)-2.14640351719968974225E-6, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)9.10010780076391431042E-8, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)-2.40274520828250956942E-9, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)3.59233385440928410398E-11, length);
        }
        Mul(t4Local, t4Local, t2Local, length);

        Muls(t5Local, t5Local, (DTYPE_X)0.0, length);
        Adds(t5Local, t5Local, (DTYPE_X)-6.31839869873368190192E-1, length);
        Add(t5Local, t5Local, t2Local, length);
        {
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)2.36706788228248691528E-1, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)-5.31806367003223277662E-2, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)8.48041718586295374409E-3, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)-9.47996768486665330168E-4, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)7.81025592944552338085E-5, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)-4.55875153252442634831E-6, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)1.89100358111421846170E-7, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)-4.91324691331920606875E-9, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)7.18466403235734541950E-11, length);
        }
        Mul(t5Local, t5Local, xLocal, length);
        Div(yLocal, t4Local, t5Local, length);
        Add(yLocal, t3Local, yLocal, length);
        Muls(yLocal, yLocal, (DTYPE_Y)0.5, length);
    }

    __aicore__ inline void _calc_condition_le_one_e_nine(LocalTensor<DTYPE_X> &xLocal, 
                                                                LocalTensor<DTYPE_Y> &yLocal,
                                                                uint32_t length)
    {
        LocalTensor<DTYPE_X> t1Local = t1Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t2Local = t2Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t3Local = t3Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t4Local = t4Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> t5Local = t5Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> all1 = All1.Get<DTYPE_X>();

        Mul(t1Local, xLocal, xLocal, length);
        Div(t2Local, all1, t1Local, length);
        Div(t3Local, all1, xLocal, length);

        Muls(t4Local, t4Local, (DTYPE_X)0.0, length);
        Adds(t4Local, t4Local, (DTYPE_X)-5.90592860534773254987E-1, length);
        {
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)6.29235242724368800674E-1, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)-1.72858975380388136411E-1, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)1.64837047825189632310E-2, length);
            Mul(t4Local, t4Local, t2Local, length);
            Adds(t4Local, t4Local, (DTYPE_X)-4.86827613020462700845E-4, length);
        }
        Mul(t4Local, t4Local, t2Local, length);

        Muls(t5Local, t5Local, (DTYPE_X)0.0, length);
        Adds(t5Local, t5Local, (DTYPE_X)-2.69820057197544900361E0, length);
        Add(t5Local, t5Local, t2Local, length);
        {
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)1.73270799045947845857E0, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)-3.93708582281939493482E-1, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)3.44278924041233391079E-2, length);
            Mul(t5Local, t5Local, t2Local, length);
            Adds(t5Local, t5Local, (DTYPE_X)-9.73655226040941223894E-4, length);
        }
        Mul(t5Local, t5Local, xLocal, length);
        Div(yLocal, t4Local, t5Local, length);
        Add(yLocal, t3Local, yLocal, length);
        Muls(yLocal, yLocal, (DTYPE_Y)0.5, length);
    }

    __aicore__ inline void _calc_condition_gt_one_e_nine(LocalTensor<DTYPE_X> &xLocal, 
                                                        LocalTensor<DTYPE_Y> &yLocal,
                                                        uint32_t length)
    {
        LocalTensor<DTYPE_X> t1Local = t1Buf.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> all1 = All1.Get<DTYPE_X>();
        Muls(t1Local, all1, (DTYPE_X)0.5, length);
        Div(yLocal, t1Local, xLocal, length);
    }

    GlobalTensor<DTYPE_X> xGm;
    GlobalTensor<DTYPE_Y> yGm;

    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> xQueue;
    TQue<QuePosition::VECOUT, BUFFER_NUM> yQueue;
    TBuf<QuePosition::VECCALC> t1Buf, t2Buf, t3Buf, t4Buf, t5Buf, rBuf;
    TBuf<QuePosition::VECCALC> Sign, Bits, Mask, Mask2;
    TBuf<QuePosition::VECCALC> All1, All0, All;

    uint32_t coreDataNum = 0;
    uint32_t tileDataNum = 0;
    uint32_t tileNum = 0;
};

// class KernelDawsn_Fp16
// {
// public:
//     __aicore__ inline KernelDawsn_Fp16() {}
//     __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,
//                             uint32_t ALIGN_NUM, uint32_t core_size, 
//                             uint32_t core_remain, uint32_t ub_data_num)
//     {
//         ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
//         uint32_t blockIdx = GetBlockIdx();

//         this->coreDataNum = core_size + (GetBlockNum() == GetBlockIdx() + 1 ? core_remain : 0);
//         this->tileDataNum = ub_data_num;

//         auto startPointer = core_size * GetBlockIdx();
//         xGm.SetGlobalBuffer((__gm__ half *)x + startPointer, this->coreDataNum);
//         yGm.SetGlobalBuffer((__gm__ half *)y + startPointer, this->coreDataNum);

//         this->tileNum = this->coreDataNum / this->tileDataNum + (this->coreDataNum % this->tileDataNum > 0);

//         pipe.InitBuffer(xQueue, BUFFER_NUM, this->tileDataNum * sizeof(half));
//         pipe.InitBuffer(yQueue, BUFFER_NUM, this->tileDataNum * sizeof(half));
//         pipe.InitBuffer(All0, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(All1, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(t1Buf, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(t2Buf, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(t3Buf, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(t4Buf, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(t5Buf, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(Sign, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(Bits, this->tileDataNum * sizeof(uint8_t));
//         pipe.InitBuffer(Mask, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(Mask2, this->tileDataNum * sizeof(float));

//         pipe.InitBuffer(All, this->tileDataNum * sizeof(float));

//         pipe.InitBuffer(xBuf, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(yBuf, this->tileDataNum * sizeof(float));
//         pipe.InitBuffer(rBuf, this->tileDataNum * sizeof(float));


//     }

//     __aicore__ inline void Process()
//     {
//         uint32_t loopCount = this->tileNum;
//         for (int32_t i = 0; i < loopCount - 1; i++) 
//         {
//             CopyIn(i, this->tileDataNum);
//             Compute(i, this->tileDataNum);
//             CopyOut(i, this->tileDataNum);
//         }

//         uint32_t length = this->coreDataNum - this->tileDataNum * (loopCount - 1);
//         CopyIn(loopCount - 1, length);
//         Compute(loopCount - 1, length);
//         CopyOut(loopCount - 1, length);
//     }

// private:

//     __aicore__ inline void CopyIn(int32_t progress, uint32_t length) 
//     {
//         LocalTensor<half> xLocal = xQueue.AllocTensor<half>();

//         DataCopy(xLocal, xGm[progress * this->tileDataNum], length);
        
//         xQueue.EnQue(xLocal);
//     }

//     __aicore__ inline void Compute(int32_t progress, uint32_t length)
//     {
//         LocalTensor<half> xLocal = xQueue.DeQue<half>();
//         LocalTensor<half> yLocal = yQueue.AllocTensor<half>();


//         LocalTensor<float> all0 = All0.Get<float>();
//         LocalTensor<float> all1 = All1.Get<float>();
//         LocalTensor<uint8_t> bits = Bits.Get<uint8_t>();
//         LocalTensor<float> sign = Sign.Get<float>();
//         LocalTensor<float> mask = Mask.Get<float>();
//         LocalTensor<float> mask2 = Mask2.Get<float>();

//         LocalTensor<float> all = All.Get<float>();


//         LocalTensor<float> xLocal_f = xBuf.Get<float>();
//         LocalTensor<float> yLocal_f = yBuf.Get<float>();
//         LocalTensor<float> rLocal = rBuf.Get<float>();


//         Cast(xLocal_f, xLocal, RoundMode::CAST_NONE, length);
        
//         Muls(all1, all1, (DTYPE_X)0.0, length);
//         Adds(all1, all1, (DTYPE_X)1.0, length);
//         Muls(all0, all0, (DTYPE_X)0.0, length);
//         Muls(rLocal, rLocal, (DTYPE_X)0.0, length);


//         Compare(bits, xLocal_f, all0, CMPMODE::GT, length);     //1 1 0 0 1 0 1
//         Select(sign, bits, all1, (DTYPE_X)-1.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         Abs(xLocal_f, xLocal_f, length);

//         // Muls(all, all, (DTYPE_X)0.0, length);
//         // Adds(all, all, (DTYPE_X)3.25, length);
//         // Compare(bits, xLocal_f, all, CMPMODE::LT, length);
//         // Select(mask, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         _calc_condition_lt_three_p_two_five(xLocal_f, yLocal_f, length);
//         // Mul(yLocal_f, yLocal_f, mask, length);
//         // Add(rLocal, rLocal, yLocal_f, length);

//         // Muls(all, all, (DTYPE_X)0.0, length);
//         // Adds(all, all, (DTYPE_X)3.25, length);
//         // Compare(bits, xLocal_f, all, CMPMODE::GE, length);
//         // Select(mask, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         // Muls(all, all, (DTYPE_X)0.0, length);
//         // Adds(all, all, (DTYPE_X)6.25, length);
//         // Compare(bits, xLocal_f, all, CMPMODE::LT, length);
//         // Select(mask2, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         // Add(mask, mask, mask2, length);
//         // Muls(all, all, (DTYPE_X)0.0, length);
//         // Adds(all, all, (DTYPE_X)2.0, length);
//         // Compare(bits, mask, all, CMPMODE::EQ, length);
//         // Select(mask2, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         // _calc_condition_lt_six_p_two_five(xLocal_f, yLocal_f, length);
//         // Mul(yLocal_f, yLocal_f, mask2, length);
//         // Add(rLocal, rLocal, yLocal_f, length);


//         // Muls(all, all, (DTYPE_X)0.0, length);
//         // Adds(all, all, (DTYPE_X)6.25, length);
//         // Compare(bits, xLocal_f, all, CMPMODE::GE, length);
//         // Select(mask, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         // Muls(all, all, (DTYPE_X)0.0, length);
//         // Adds(all, all, (DTYPE_X)1e9, length);
//         // Compare(bits, xLocal_f, all, CMPMODE::LE, length);
//         // Select(mask2, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         // Add(mask, mask, mask2, length);
//         // Muls(all, all, (DTYPE_X)0.0, length);
//         // Adds(all, all, (DTYPE_X)2.0, length);
//         // Compare(bits, mask, all, CMPMODE::GE, length);
//         // Select(mask2, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         // _calc_condition_le_one_e_nine(xLocal_f, yLocal_f, length);
//         // Mul(yLocal_f, yLocal_f, mask2, length);
//         // Add(rLocal, rLocal, yLocal_f, length);

//         // Muls(all, all, (DTYPE_X)0.0, length);
//         // Adds(all, all, (DTYPE_X)1e9, length);
//         // Compare(bits, xLocal_f, all, CMPMODE::GT, length);
//         // Select(mask, bits, all1, (DTYPE_X)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, length);
//         // _calc_condition_gt_one_e_nine(xLocal_f, yLocal_f, length);
//         // Mul(yLocal_f, yLocal_f, mask, length);
//         // Add(rLocal, rLocal, yLocal_f, length);
        
//         Mul(yLocal_f, rLocal, sign, length);
        
//         Cast(yLocal, yLocal_f, RoundMode::CAST_NONE, length);

//         yQueue.EnQue(yLocal);
//         xQueue.FreeTensor(xLocal);
//     }

//     __aicore__ inline void CopyOut(int32_t progress, uint32_t length) {
//         LocalTensor<half> yLocal = yQueue.DeQue<half>();
//         DataCopy(yGm[progress * this->tileDataNum], yLocal, length);
//         yQueue.FreeTensor(yLocal);
//     }

//     __aicore__ inline void _calc_condition_lt_three_p_two_five(LocalTensor<float> &xLocal, 
//                                                                 LocalTensor<float> &yLocal,
//                                                                 uint32_t length)
//     {
//         LocalTensor<float> t1Local = t1Buf.Get<float>();
//         LocalTensor<float> t2Local = t2Buf.Get<float>();
//         LocalTensor<float> t3Local = t3Buf.Get<float>();


//         Mul(t1Local, xLocal, xLocal, length);

//         Muls(t2Local, t2Local, (float)0.0, length);
//         Adds(t2Local, t2Local, (float)1.13681498971755972054E-11, length);

//         {
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)8.49262267667473811108E-10, length);
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)1.94434204175553054283E-8, length);
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)9.53151741254484363489E-7, length);
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)3.07828309874913200438E-6, length);
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)3.52513368520288738649E-4, length);
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)-8.50149846724410912031E-4, length);
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)4.22618223005546594270E-2, length);
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)-9.17480371773452345351E-2, length);
//             Mul(t2Local, t1Local, t2Local, length);
//             Adds(t2Local, t2Local, (float)9.99999999999999994612E-1, length);
//         }

//         Muls(t3Local, t3Local, (float)0.0, length);
//         Adds(t3Local, t3Local, (float)1.13681498971755972054E-11, length);
        
//         {
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)1.48864681368493396752E-9, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)5.21265281010541664570E-8, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)1.27258478273186970203E-6, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)2.32490249820789513991E-5, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)3.25524741826057911661E-4, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)3.48805814657162590916E-3, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)2.79448531198828973716E-2, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)1.58874241960120565368E-1, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)5.74918629489320327824E-1, length);
//             Mul(t3Local, t1Local, t3Local, length);
//             Adds(t3Local, t3Local, (float)1.00000000000000000539E0, length);
//         }

//         Mul(yLocal, xLocal, t2Local, length);
//         Div(yLocal, yLocal, t3Local, length);
//     }

//     __aicore__ inline void _calc_condition_lt_six_p_two_five(LocalTensor<float> &xLocal, 
//                                                                 LocalTensor<float> &yLocal,
//                                                                 uint32_t length)
//     {
//         LocalTensor<float> t1Local = t1Buf.Get<float>();
//         LocalTensor<float> t2Local = t2Buf.Get<float>();
//         LocalTensor<float> t3Local = t3Buf.Get<float>();
//         LocalTensor<float> t4Local = t4Buf.Get<float>();
//         LocalTensor<float> t5Local = t5Buf.Get<float>();
//         LocalTensor<float> all1 = All1.Get<float>();

//         Mul(t1Local, xLocal, xLocal, length);
//         Div(t2Local, all1, t1Local, length);
//         Div(t3Local, all1, xLocal, length);


//         Muls(t4Local, t4Local, (float)0.0, length);
//         Adds(t4Local, t4Local, (float)5.08955156417900903354E-1, length);
//         {
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)-2.44754418142697847934E-1, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)9.41512335303534411857E-2, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)-2.18711255142039025206E-2, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)3.66207612329569181322E-3, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)-4.23209114460388756528E-4, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)3.59641304793896631888E-5, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)-2.14640351719968974225E-6, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)9.10010780076391431042E-8, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)-2.40274520828250956942E-9, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)3.59233385440928410398E-11, length);
//         }
//         Mul(t4Local, t4Local, t2Local, length);

//         Muls(t5Local, t5Local, (float)0.0, length);
//         Adds(t5Local, t5Local, (float)-6.31839869873368190192E-1, length);
//         Add(t5Local, t5Local, t2Local, length);
//         {
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)2.36706788228248691528E-1, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)-5.31806367003223277662E-2, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)8.48041718586295374409E-3, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)-9.47996768486665330168E-4, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)7.81025592944552338085E-5, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)-4.55875153252442634831E-6, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)1.89100358111421846170E-7, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)-4.91324691331920606875E-9, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)7.18466403235734541950E-11, length);
//         }
//         Mul(t5Local, t5Local, xLocal, length);
//         Div(yLocal, t4Local, t5Local, length);
//         Add(yLocal, t3Local, yLocal, length);
//         Muls(yLocal, yLocal, (float)0.5, length);
//     }

//     __aicore__ inline void _calc_condition_le_one_e_nine(LocalTensor<float> &xLocal, 
//                                                                 LocalTensor<float> &yLocal,
//                                                                 uint32_t length)
//     {
//         LocalTensor<float> t1Local = t1Buf.Get<float>();
//         LocalTensor<float> t2Local = t2Buf.Get<float>();
//         LocalTensor<float> t3Local = t3Buf.Get<float>();
//         LocalTensor<float> t4Local = t4Buf.Get<float>();
//         LocalTensor<float> t5Local = t5Buf.Get<float>();
//         LocalTensor<float> all1 = All1.Get<float>();

//         Mul(t1Local, xLocal, xLocal, length);
//         Div(t2Local, all1, t1Local, length);
//         Div(t3Local, all1, xLocal, length);

//         Muls(t4Local, t4Local, (float)0.0, length);
//         Adds(t4Local, t4Local, (float)-5.90592860534773254987E-1, length);
//         {
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)6.29235242724368800674E-1, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)-1.72858975380388136411E-1, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)1.64837047825189632310E-2, length);
//             Mul(t4Local, t4Local, t2Local, length);
//             Adds(t4Local, t4Local, (float)-4.86827613020462700845E-4, length);
//         }
//         Mul(t4Local, t4Local, t2Local, length);

//         Muls(t5Local, t5Local, (float)0.0, length);
//         Adds(t5Local, t5Local, (float)-2.69820057197544900361E0, length);
//         Add(t5Local, t5Local, t2Local, length);
//         {
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)1.73270799045947845857E0, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)-3.93708582281939493482E-1, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)3.44278924041233391079E-2, length);
//             Mul(t5Local, t5Local, t2Local, length);
//             Adds(t5Local, t5Local, (float)-9.73655226040941223894E-4, length);
//         }
//         Mul(t5Local, t5Local, xLocal, length);
//         Div(yLocal, t4Local, t5Local, length);
//         Add(yLocal, t3Local, yLocal, length);
//         Muls(yLocal, yLocal, (float)0.5, length);
//     }

//     GlobalTensor<half> xGm;
//     GlobalTensor<half> yGm;

//     TPipe pipe;
//     TQue<QuePosition::VECIN, BUFFER_NUM> xQueue;
//     TQue<QuePosition::VECOUT, BUFFER_NUM> yQueue;
//     TBuf<QuePosition::VECCALC> t1Buf, t2Buf, t3Buf, t4Buf, t5Buf, rBuf;
//     TBuf<QuePosition::VECCALC> Sign, Bits, Mask, Mask2;
//     TBuf<QuePosition::VECCALC> All1, All0, All;
//     TBuf<QuePosition::VECCALC> xBuf, yBuf;

//     uint32_t coreDataNum = 0;
//     uint32_t tileDataNum = 0;
//     uint32_t tileNum = 0;
// };

extern "C" __global__ __aicore__ void dawsn(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    // TODO: user kernel impl
    // if (tiling_data.ALIGN_NUM == 16)
    // {
    //     KernelDawsn_Fp16 op;
    //     op.Init(x, y, 
    //             tiling_data.ALIGN_NUM, tiling_data.core_size, 
    //             tiling_data.core_remain, tiling_data.ub_data_num);
    //     op.Process();
    // }
    // else
    // {
        KernelDawsn op;
        op.Init(x, y, 
                tiling_data.ALIGN_NUM, tiling_data.core_size, 
                tiling_data.core_remain, tiling_data.ub_data_num);
        op.Process();
    // }
}